You write custom CUDA kernels to replace pytorch operators in given architecture to get speedups. You have complete freedom to choose set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into a single kernel, for example, combining matmul+relu), or algorithmic changes (such as online softmax). You are only limited by your imagination.

**SPECIAL INSTRUCTIONS FOR CHEBYSHEV + LEAKYRELU FUSION:**

When implementing Chebyshev Distance + LeakyReLU fusion, you MUST implement the following optimized strategy:

1. **FUSION ARCHITECTURE**: Combine LeakyReLU activation and Chebyshev distance computation in a single kernel:
   - Apply LeakyReLU activation to input tensor x first
   - Compute Chebyshev distance between activated x and y
   - Eliminate intermediate tensor storage for maximum efficiency

2. **WARP-LEVEL OPTIMIZATION**: Use warp-level processing for maximum performance:
   - Each block processes one sample from the batch
   - Use 8 warps per block (256 threads) for optimal GPU utilization
   - Use __shfl_down_sync for efficient warp-level reduction of maximum values
   - Divide feature dimensions among warps for parallel processing

3. **MEMORY COALESCING**: Ensure efficient memory access patterns:
   - Each thread processes multiple elements with stride pattern
   - Store LeakyReLU results directly to global memory
   - Use coalesced memory access for both input tensors

4. **EFFICIENT MAXIMUM REDUCTION**: Implement optimized maximum value finding:
cpp
// Warp-level maximum reduction
for (int offset = 16; offset > 0; offset /= 2) {
    float other_max = __shfl_down_sync(0xffffffff, warp_max, offset);
    if (other_max > warp_max) {
        warp_max = other_max;
    }
}

// Cross-warp reduction using shared memory
extern __shared__ float shared_data[];
if (lane_id == 0) {
    shared_data[warp_id] = warp_max;
}
__syncthreads();



5. **LEAKYRELU FUSION**: Integrate LeakyReLU activation seamlessly:
cpp
// Apply LeakyReLU activation while computing distance
float x_activated = x[base + dim] > 0.0f ? x[base + dim] : negative_slope * x[base + dim];

// Store activated output
leakyrelu_output[base + dim] = x_activated;

// Compute Chebyshev distance
float diff = x_activated - y[base + dim];
float abs_diff = fabsf(diff);
if (abs_diff > warp_max) {
    warp_max = abs_diff;
}



6. **SHARED MEMORY PATTERN**: Use efficient shared memory organization:
cpp
// For maximum reduction across warps
extern __shared__ float shared_data[];
if (lane_id == 0) {
    shared_data[warp_id] = warp_max;
}
__syncthreads();

// Final maximum calculation
if (tid == 0) {
    float global_max = 0.0f;
    int num_warps = blockDim.x / 32;
    for (int i = 0; i < num_warps; i++) {
        if (shared_data[i] > global_max) {
            global_max = shared_data[i];
        }
    }
    distances[sample_idx] = global_max;
}



7. **BLOCK CONFIGURATION**: Use optimal settings:
   - Block size: 256 threads (8 warps)
   - Shared memory: 8 * sizeof(float) for warp reduction results
   - One block per sample for maximum parallelism
   - Elements per warp: (feature_dim + 8 - 1) / 8

8. **PRECISION REQUIREMENTS**: Ensure exact mathematical alignment:
   - LeakyReLU: activated_x = x > 0 ? x : negative_slope * x
   - Chebyshev Distance: max(|activated_x - y|)
   - Use fabsf for absolute value computation
   - Verify with torch.allclose(rtol=1e-03, atol=1e-6)

9. **FUNCTION SIGNATURE**: The main CUDA function must accept all parameters:
cpp
torch::Tensor chebyshev_leakyrelu_cuda(
    torch::Tensor x,
    torch::Tensor y,
    float negative_slope
)



10. **MATHEMATICAL FORMULAS**: Implement exact mathematical operations:
    - LeakyReLU Activation: leakyrelu(x) = max(x, negative_slope * x)
    - Absolute Difference: abs_diff = |leakyrelu(x) - y|
    - Chebyshev Distance: chebyshev_dist = max(abs_diff)

11. **PYTHON CALLING CONVENTION**: The ModelNew forward method must pass parameters correctly:
python
def forward(self, x, y):
    return self.chebyshev_leakyrelu.chebyshev_leakyrelu_cuda(x, y, self.negative_slope)



12. **OUTPUT REQUIREMENTS**: Generate both distances and activated outputs:
    - Primary output: Chebyshev distances [batch_size]
    - Secondary output: LeakyReLU activated tensor [batch_size, feature_dim]
    - Both outputs must match PyTorch reference implementation exactly

13. **PERFORMANCE OPTIMIZATIONS**: Include advanced optimizations:
    - Use fast math optimizations (--use_fast_math)
    - Optimize for compute capability 8.0+ (sm_80)
    - Use -O3 optimization level
    - Avoid bank conflicts in shared memory access

Here's the target architecture to optimize:

python
import torch
import torch.nn as nn

class Model(nn.Module):
"""
Chebyshev Distance implementation.
Computes the Chebyshev distance (maximum absolute difference) between two sets of vectors.
"""
def init(self):
super(Model, self).init()

def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    """
    Compute Chebyshev distance between x and y.

    Args:
        x (torch.Tensor): First set of vectors [batch_size, feature_dim]
        y (torch.Tensor): Second set of vectors [batch_size, feature_dim]

    Returns:
        torch.Tensor: Chebyshev distances [batch_size]
    """
    # Compute absolute differences
    abs_diff = torch.abs(x - y)
    
    # Find maximum along feature dimension
    chebyshev_dist = torch.max(abs_diff, dim=1)[0]
    
    return chebyshev_dist

batch_size = 256
feature_dim = 512

def get_inputs():
# Generate two sets of vectors
x = torch.randn(batch_size, feature_dim)
y = torch.randn(batch_size, feature_dim)
return [x, y]

def get_init_inputs():
return [] # No special initialization inputs needed



**EXPECTED OUTPUT STRUCTURE**:
Generate two files:
1. `chebyshev_leakyrelu_cudacode.py` - Contains ModelNew class with Chebyshev+LeakyReLU fusion using pure CUDA
2. `chebyshev_leakyrelu_torchcode.py` - Contains the reference PyTorch implementation with LeakyReLU fusion

**KEY REQUIREMENTS**:
- The CUDA implementation must use pure CUDA functions only
- Must implement LeakyReLU activation before Chebyshev distance computation
- Must use warp-level optimization for maximum performance
- Must use efficient maximum reduction algorithm
- Must handle arbitrary tensor shapes (not just fixed dimensions)
- Must maintain mathematical precision with PyTorch implementation
- Must use optimal block configuration (256 threads, 8 warps)
- Expected speedup: 1.8-2.5x over PyTorch baseline
- Must use fast math optimizations for better performance
- Must be robust and handle edge cases properly
- Must use only pure CUDA functions (no PyTorch internal functions)
- Must use fabsf for absolute value computation
- Must implement exact mathematical formulas for LeakyReLU and Chebyshev distance
- Must pass negative_slope parameter correctly from Python to CUDA
- Must use Python float syntax (0.01) not C++ syntax (0.01f) in Python code
- Must generate both distance and activated output tensors
- Must use shared memory efficiently for warp-level maximum reduction
- Must ensure coalesced memory access patterns
- Must eliminate intermediate tensor storage for maximum fusion benefits
